refactor(rocm): remove the gpu_iface layer and flatten the ROCm headers - #329

Applies the rocm/-subfolder convention to the C++ tree: the ROCm sources
become flashinfer/csrc/rocm/ rather than a _rocm suffix on the parent.

Deliberately package-internal rather than the repo-root csrc/rocm/ the plan
named. Root csrc/ is upstream's and now 568 files at v0.6.18rc3, so putting
fork-only sources inside it invites incidental conflict for no gain -- the
directory has no upstream counterpart either way, so neither location
changes the merge surface. It would also cost a second build-backend
materialisation step: root csrc/ is outside the package, so the wheel would
need include/'s copy-into-package treatment, adding a path where a wheel
ships without the sources its JIT needs.

flashinfer/csrc/ does not exist upstream today, so there is no add/add
conflict now. It is a plainer name than a _rocm suffix would be, so this
accepts a small future add/add risk against reading better; if upstream ever
adds flashinfer/csrc/, the resolution is to keep both subtrees.

get_csrc_dir(), the package-data globs and the MANIFEST graft move with it;
51 references across 16 files follow. CHANGELOG.md keeps the old name --
those entries describe releases that shipped it.
docs/rocm/coverage-gfx942.json also keeps it: that file is a measurement
snapshot taken at 0e68a2f, where csrc_rocm/ was the correct path, and it
is regenerated by a gfx942 coverage run rather than edited by hand.

Verified by diffing the full wheel inventory, base vs HEAD, rather than
spot-checking: 437 -> 438 entries, the 65 csrc_rocm/ paths replaced
one-for-one by csrc/rocm/, and the single addition is
flashinfer/rocm/torch_compile.py from the previous commit. Nothing dropped.
The 65 includes aiter_tensor_compat.h, which #294 added under the old path
after this branch started -- the rebase resolution is a rename, and `git
rebase --skip` would have dropped it and broken three *_aiter.cu compiles.

test_build_backend.py::test_wheel_carries_the_paths_the_jit_resolves asserts
the same property from inside the suite and passes once `build` is installed;
it is importorskip'd, so it is silently absent in an environment without it.

gpu_iface/ and rocm/attention/ each carried a dispatch.cuh and an
exception.h. Both dispatch copies define DISPATCH_NUM_MMA_KV with different
bodies and both use #pragma once, so a TU reaching both silently redefines
the macro:

  attention/dispatch.cuh:31:9: warning: 'DISPATCH_NUM_MMA_KV' macro redefined

Nothing surfaced it because the build does not pass -Wmacro-redefined.

Merged to include/flashinfer/rocm/{dispatch.cuh,exception.h}, above
attention/ because rocm/sampling.cuh needs them too.

dispatch.cuh keeps attention's DISPATCH_NUM_MMA_KV -- the one that falls
back to NUM_MMA_KV=1 rather than raising -- and gains gpu_iface's
DISPATCH_ROPE_DIM, which had no counterpart. That is behaviour-preserving,
not a preference: all seven expansion sites live in prefill.cuh (1694, 2485,
2600), pod.cuh (234, 246) and batch_pod.cuh (212, 223), and each of those
three includes cascade.cuh -- which pulls gpu_iface's copy -- immediately
before the sibling "dispatch.cuh". Attention's definition therefore already
won at every site. gpu_iface's raising body was defined in other TUs but
reached no expansion site, so deleting it changes nothing that ran.

exception.h takes gpu_iface's copy, which is the superset: it carries
#include <string>, which attention's omits while still using std::string.

The relative includes had to be qualified -- attention/{prefill,pod,
batch_pod}.cuh said "dispatch.cuh" and attention/{allocator.h,scheduler.cuh}
said "exception.h", both of which resolved to the now-deleted siblings. A
missed one hard-errors rather than silently picking up the wrong file.

Verified on gfx942 and gfx950: the redefinition probe goes from 1 warning to
0, and cascade, scheduler, norm, pos_enc, page, sampling and quantization
each compile -fsyntax-only with zero errors and zero redefinitions.

rocm/sampling.cuh included upstream's <flashinfer/allocator.h>, which
includes upstream's exception.h. Our exception.h claims the same guard,
FLASHINFER_EXCEPTION_H_, and wins the race -- so upstream's compiled to
nothing and its allocator was built against our header instead. Preprocessing
rocm/sampling.cuh showed our copy expanded first and upstream's producing no
output; grep for 'class Warning' and FLASHINFER_WARN over the result returned
0, confirming upstream's declarations were silently absent. It worked only
because allocator.h happens not to use them.

Point it at the fork's rocm/attention/allocator.h, which differs from
upstream's only in the overflow diagnostic and provides a superset of what
sampling.cuh uses (AlignedAllocator, aligned_alloc<T>). The fork copy's own
include closure is just itself plus rocm/exception.h, so nothing upstream
comes back in through the side door.

Also drops the #else arm that selected the CUDA headers. gpu_iface/macros.hpp
#errors on a non-HIP compiler, so that branch could never be taken; it was
the only place a ROCm header named <flashinfer/math.cuh> and
<flashinfer/vec_dtypes.cuh>, both of which share guards with our HIP copies.

The closure of rocm/sampling.cuh is now entirely fork-owned -- flashinfer/rocm/*
and gpu_iface/* only, zero include/flashinfer/*.h -- which is what makes the
tripwires in a later commit meaningful rather than a list of exemptions.
Compiles -fsyntax-only with 0 errors and 0 redefinitions on gfx942 and gfx950.

memory_ops_hip.h was not standalone-compilable. It names PrefetchMode and
SharedMemFillMode unqualified and got them only from the scope its includer
created: memory_ops.hpp included it from *inside*
flashinfer::gpu_iface::memory, so the backend's own namespace chain nested
inside that one and the real FQN was doubled --
flashinfer::gpu_iface::memory::flashinfer::gpu_iface::memory::detail::hip.
The alias on the next line resolved only because unqualified lookup finds the
outer `flashinfer` again from in there.

Extracting the two enums to memory_types.hpp lets the backend include what it
needs and the wrapper include it at global scope. Proof that this is a real
fix rather than a textual one: a TU naming
flashinfer::gpu_iface::memory::detail::hip::commit_group fails before with

  error: no member named 'detail' in namespace 'flashinfer::gpu_iface::memory'

and compiles after. Rewriting only the alias would have left it compiling via
the same lookup trick, with the doubling intact.

vec_dtypes_hip.h had the same shape but no namespace of its own at all -- it
opened a bare `detail::hip`, so lifting its include to global scope would have
put vec_t and vec_cast in the global namespace. It now declares
flashinfer::gpu_iface::vec_dtypes explicitly, opened *after* its system
includes: those were previously expanded inside three fork namespaces and
survived only because platform.hpp had already pulled them at global scope.
Its anonymous namespace, which shadows HIP's own __float2bfloat162_rn, stays
nested where it is today rather than becoming ambiguous at global scope.

Also deletes error.hpp and backend/hip/mma_debug_utils_hip.h -- zero
includers; the only references to gpu_iface::debug_utils were that header's
own namespace braces.

All 45 fork headers compile -fsyntax-only standalone with zero errors and
zero macro redefinitions on both gfx942 and gfx950.

Since #320 deleted the CUDA backend, include/gpu_iface/ has been HIP-only
headers sitting at the include root as if it were a shared abstraction, with a
backend/hip/ level implying a backend selection that does not exist --
macros.hpp #errors on a non-HIP compiler.

19 headers move up into include/flashinfer/rocm/, losing both directory
levels. No basename collisions: the four backend files all carry _hip
suffixes, and the two that would have clashed with rocm/attention/ were merged
in an earlier commit. The _hip suffixes stay -- with wrapper and backend now
in one directory, that suffix is what distinguishes them.

~140 include lines follow, including 22 across 14 .jinja templates and one in
a Python string at jit/rocm/activation.py, neither of which any C++ tool would
have caught.

upstream_canary pairs forked headers to upstream by basename, so the flatten
puts exception.h, fastdiv.cuh and layout.cuh into the drift report for the
first time -- all three are genuine forks and were previously invisible to it.
utils.cuh is excluded: it shares upstream's basename but 411 of 443 lines
differ, so it is a HIP rewrite that never tracked upstream and would report
drift forever. The exclusion is an exact-path set rather than a directory
prefix, with a test that a sibling is still paired.

Verified on gfx942 and gfx950: all 45 fork headers compile -fsyntax-only
standalone with zero errors and zero macro redefinitions. Negative control: a
TU including the old <gpu_iface/platform.hpp> now fails "file not found", so
nothing is still resolving the old layout from another include root.

gpu_iface described a backend-selection layer that no longer exists: #320
deleted the CUDA half, and macros.hpp #errors on a non-HIP compiler. What
remained was five levels of namespace -- flashinfer::gpu_iface::vec_dtypes::
detail::hip::vec_t -- for what upstream calls flashinfer::vec_t.

Levels are kept only where upstream also groups: math and memory survive
(upstream has flashinfer::math and flashinfer::cp_async). gpu_iface, the
single-backend vec_dtypes wrapper and the detail::hip impl split all go, since
there is no second backend to disambiguate against.

mma is the exception. Its MMAMode, mma_sync_m16n16k16_row_col_f16f16f32 and
m16k16_rowsum_f16f16f32 have exact-signature counterparts in upstream's
flashinfer::mma (mma.cuh:57,318,522), and both headers use bare #pragma once,
so no include guard would catch the collision. They land in
flashinfer::mma_hip instead.

Four headers turned out to be pure pass-throughs once the namespace went --
math_ops, vec_dtypes, memory_ops and mma_ops each re-exported the backend into
the namespace the backend now declares itself, which is a redefinition rather
than a shim. They are deleted and their includers point at the backend
headers, which absorbed the macros.hpp and platform.hpp includes the wrappers
had been contributing.

conversion_utils.h's `namespace fi::con` was the only namespace in the repo
not rooted at flashinfer; explicit_casting joins flashinfer.

Two things no C++ tool would have found: 22 include lines across 14 .jinja
templates plus jit/rocm/activation.py, and five gpu_iface::math uses inside
the JIT-compiled C++ string literal in
test_customize_prefill_use_softmax_hip.py, which fails at test time rather
than build time.

`grep -rn gpu_iface` now returns only CHANGELOG history plus the docs and the
pre-commit path, both retargeted in the next commit. All 42 fork headers
compile -fsyntax-only standalone with zero errors and zero macro redefinitions
on gfx942 and gfx950.

The tree diagram, the intrinsics paragraph and the add-rocm-kernel skill all
described a gpu_iface/backend/hip/ layout and a namespace that no longer
exist. The paragraph's instruction -- "add the abstraction in gpu_iface/ and
implement it under gpu_iface/backend/hip/" -- would have been actively
misleading.

Also rewrites the naming rule. It said to prefer a _rocm/_aiter suffix for
anything upstream might plausibly create; the port has since settled on a
rocm/ subdirectory instead, which is what flashinfer/csrc/rocm/ and
include/flashinfer/rocm/ are, and which collides with nothing even as upstream
grows those trees. The suffix stays as the fallback where a subdirectory does
not fit, which is now just the _hip.h intrinsic headers.

The forked-headers paragraph gains the three headers the flatten newly exposed
to the drift report (layout.cuh, fastdiv.cuh, exception.h) and names utils.cuh
as the one exclusion, so the canary's exclude set is documented where a reader
looks for it rather than only in the script.

The reuse pre-commit pattern loses its dead include/gpu_iface/ alternative and
gains include/flashinfer/attention/aiter/ -- six AMD-authored headers that no
licence check has ever reached.

CLAUDE.md needed no change; it never mentioned gpu_iface.

22 fork headers reused their upstream namesake's include guard, so whichever
arrived first silently suppressed the other. That is how the sampling path
came to compile upstream's allocator.h against our exception.h, fixed two
commits ago -- the failure mode is a wrong-but-compatible header vanishing,
which compiles.

Each now defines FLASHINFER_ROCM_<PATH>_ instead, plus an #error if the
upstream guard is already set. Both include orders are now loud:

  upstream first -> "include/flashinfer/exception.h and
    include/flashinfer/rocm/exception.h both define FLASHINFER_EXCEPTION_H_;
    include only one"
  fork first     -> upstream's header is no longer suppressed, so it reaches
                    the compiler and produces redefinition errors

Ugly in the second direction, but ugly-and-loud beats silent. Upstream's
headers are not touched; the additive-only policy rules that out, which is why
only one direction gets the sentence.

Safe to do now and not before: a preprocessor scan of all 42 fork headers
shows none reaches an upstream include/flashinfer/*.h any more, so nothing was
relying on the suppression. It matters more after the flatten, because
include/flashinfer/rocm/utils.cuh now reads like a peer of
include/flashinfer/utils.cuh, whose ten same-signature functions it shadows.

Note upstream math.cuh and vec_dtypes.cuh cannot be reached from a ROCm TU
regardless -- they include cuda_fp16.h -- so their guards were never the live
risk. exception.h, allocator.h, utils.cuh, layout.cuh and fastdiv.cuh are
pure C++ and were.

All 42 headers compile -fsyntax-only with zero errors and zero redefinitions
on gfx942 and gfx950; pre-commit clean.

The re-guard in the previous commit keyed on `^#ifndef`, so it silently passed
over every fork header that used a bare `#pragma once` -- including
rocm/utils.cuh, which that commit's own message names as the motivating case.
Six headers with an upstream namesake were left without a tripwire:
utils.cuh and attention/{prefill,cascade,page,pod,batch_pod}.cuh.

They never shared a guard, so neither include order was ever silent; they just
produced unnamed redefinition errors where the other 22 name the pair. They
now carry the same #ifdef/#error, so the scheme is uniform.

Also corrects three things this PR invalidated on its own:

- utils.cuh and sampling.cuh still explained an include choice by the
  FLASHINFER_EXCEPTION_H_ race, which the re-guard removed two commits later.
- sampling.cuh had an empty `#ifdef PLATFORM_HIP_DEVICE` / `#endif` left where
  a `using namespace` line was deleted.
- CONTRIBUTING and the add-rocm-kernel skill claimed upstream declares
  `flashinfer::memory` (it declares `cp_async`) and that upstream's mma.cuh has
  no include guard (it has FLASHINFER_MMA_CUH_). The real rule is that a fork
  header keeps its name and tripwires when the two are meant to coexist, and is
  renamed when they are meant to be usable together, which is why mma_hip is
  the one rename. The intrinsics rule also read as an absolute ban on hipcub
  and inline asm outside the shared headers, which six live fork headers break.

Known and accepted: the anonymous namespace in vec_dtypes_hip.h shadows HIP's
__float2bfloat162_rn, and the flatten widened that shadow from
flashinfer::gpu_iface::vec_dtypes to flashinfer. Its callers moved with it, so
in-file resolution is unchanged, and no other header calls either helper.

Three follow-ups from reviewing the flatten.

macros.hpp is what #errors on a non-HIP compiler, so it has to be included
before any hip/ header or the failure is "hip/hip_fp16.h: No such file or
directory" instead of the named diagnostic. cooperative_groups.h already
carries a clang-format-off block saying exactly that; math_hip.h and
sampling.cuh gained macros.hpp during the flatten and clang-format sorted it
after the system includes, silently defeating it. Both now use the same pin.
sampling.cuh's comment asserted the ordering it did not have.

rocm/dispatch.cuh defines eleven DISPATCH_* macros that upstream's utils.cuh
also defines. The guard tripwires added earlier only cover headers that share
an include guard, and these two share neither a guard nor a basename, so the
overlap had no cover at all -- a TU reaching both would silently take
whichever came last, flipping DISPATCH_NUM_MMA_KV between the AMD fallback and
upstream's throwing arm. It now tripwires on FLASHINFER_UTILS_CUH_.

The fork's AlignedAllocator reported only the name, size and alignment on
overflow. Pointing sampling.cuh at it (two commits ago) therefore downgraded
what a user sees when they under-size the workspace, since upstream's copy
names the shortfall and the remedy. Restores that wording, which improves
every ROCm Plan() path rather than just sampling.

vec_dtypes_hip.h defined FLASHINFER_INLINE twice, identically and 21 lines
apart. Pre-existing -- the pre-flatten copy has both -- but a PR whose subject
is removing macro redefinitions should not leave one behind.

docs/rocm/coverage-gfx942.json still named csrc_rocm/ in five `reason`
strings. Those are labels echoed from scripts/coverage_ownership.toml, which
this branch already updated, so leaving them meant the two files disagreed and
the repo carried paths that no longer exist. Only the strings change: the
execution_percent, the base sha and every per-file count are byte-identical,
asserted rather than eyeballed. The snapshot itself stays stale -- its base is
340 commits back and regenerating it needs a gfx942 coverage run -- which is a
pre-existing gap this branch does not close.